#include <iostream>
using namespace std;

__global__ void add(float* x, float* y, float* z, int n) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride) {
        z[i] = x[i] + y[i];
    }
}

int main() {
    int N = 1 << 20;
    int nBytes = N * sizeof(float);

    float *x, *y, *z;

    cudaMallocManaged((void**)&x, nBytes);
    cudaMallocManaged((void**)&y, nBytes);
    cudaMallocManaged((void**)&z, nBytes);

    for (int i = 0; i < N; i++) {
        x[i] = 10.0;
        y[i] = 20.0;
    }

    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    add <<<gridSize, blockSize>>>(x, y, z, N);

    cudaDeviceSynchronize();
    float maxError = 0.0;

    for (int i = 0; i < N; i++) {
        maxError = fmax(maxError, fabs(z[i] - 20.0));
    }

    cout << maxError << endl;

    cudaFree(x);
    cudaFree(y);
    cudaFree(z);

    return 0;
}
